// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com>
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.

#ifndef EIGEN_CXX11_TENSOR_TENSOR_STRIDING_H
#define EIGEN_CXX11_TENSOR_TENSOR_STRIDING_H

namespace Eigen {

/** \class TensorStriding
 * \ingroup CXX11_Tensor_Module
 *
 * \brief Tensor striding class.
 *
 *
 */
namespace internal {
template<typename Strides, typename XprType>
struct traits<TensorStridingOp<Strides, XprType>> : public traits<XprType>
{
	typedef typename XprType::Scalar Scalar;
	typedef traits<XprType> XprTraits;
	typedef typename XprTraits::StorageKind StorageKind;
	typedef typename XprTraits::Index Index;
	typedef typename XprType::Nested Nested;
	typedef typename remove_reference<Nested>::type _Nested;
	static const int NumDimensions = XprTraits::NumDimensions;
	static const int Layout = XprTraits::Layout;
	typedef typename XprTraits::PointerType PointerType;
};

template<typename Strides, typename XprType>
struct eval<TensorStridingOp<Strides, XprType>, Eigen::Dense>
{
	typedef const TensorStridingOp<Strides, XprType> EIGEN_DEVICE_REF type;
};

template<typename Strides, typename XprType>
struct nested<TensorStridingOp<Strides, XprType>, 1, typename eval<TensorStridingOp<Strides, XprType>>::type>
{
	typedef TensorStridingOp<Strides, XprType> type;
};

} // end namespace internal

template<typename Strides, typename XprType>
class TensorStridingOp : public TensorBase<TensorStridingOp<Strides, XprType>>
{
  public:
	typedef TensorBase<TensorStridingOp<Strides, XprType>> Base;
	typedef typename Eigen::internal::traits<TensorStridingOp>::Scalar Scalar;
	typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename Eigen::internal::nested<TensorStridingOp>::type Nested;
	typedef typename Eigen::internal::traits<TensorStridingOp>::StorageKind StorageKind;
	typedef typename Eigen::internal::traits<TensorStridingOp>::Index Index;

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorStridingOp(const XprType& expr, const Strides& dims)
		: m_xpr(expr)
		, m_dims(dims)
	{
	}

	EIGEN_DEVICE_FUNC
	const Strides& strides() const { return m_dims; }

	EIGEN_DEVICE_FUNC
	const typename internal::remove_all<typename XprType::Nested>::type& expression() const { return m_xpr; }

	EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorStridingOp)

  protected:
	typename XprType::Nested m_xpr;
	const Strides m_dims;
};

// Eval as rvalue
template<typename Strides, typename ArgType, typename Device>
struct TensorEvaluator<const TensorStridingOp<Strides, ArgType>, Device>
{
	typedef TensorStridingOp<Strides, ArgType> XprType;
	typedef typename XprType::Index Index;
	static const int NumDims = internal::array_size<typename TensorEvaluator<ArgType, Device>::Dimensions>::value;
	typedef DSizes<Index, NumDims> Dimensions;
	typedef typename XprType::Scalar Scalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
	static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
	typedef StorageMemory<CoeffReturnType, Device> Storage;
	typedef typename Storage::Type EvaluatorPointerType;

	enum
	{
		IsAligned = /*TensorEvaluator<ArgType, Device>::IsAligned*/ false,
		PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
		BlockAccess = false,
		PreferBlockAccess = TensorEvaluator<ArgType, Device>::PreferBlockAccess,
		Layout = TensorEvaluator<ArgType, Device>::Layout,
		CoordAccess = false, // to be implemented
		RawAccess = false
	};

	//===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
	typedef internal::TensorBlockNotImplemented TensorBlock;
	//===--------------------------------------------------------------------===//

	EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
		: m_impl(op.expression(), device)
	{
		m_dimensions = m_impl.dimensions();
		for (int i = 0; i < NumDims; ++i) {
			m_dimensions[i] = Eigen::numext::ceil(static_cast<float>(m_dimensions[i]) / op.strides()[i]);
		}

		const typename TensorEvaluator<ArgType, Device>::Dimensions& input_dims = m_impl.dimensions();
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			m_outputStrides[0] = 1;
			m_inputStrides[0] = 1;
			for (int i = 1; i < NumDims; ++i) {
				m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1];
				m_inputStrides[i] = m_inputStrides[i - 1] * input_dims[i - 1];
				m_inputStrides[i - 1] *= op.strides()[i - 1];
			}
			m_inputStrides[NumDims - 1] *= op.strides()[NumDims - 1];
		} else { // RowMajor
			m_outputStrides[NumDims - 1] = 1;
			m_inputStrides[NumDims - 1] = 1;
			for (int i = NumDims - 2; i >= 0; --i) {
				m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1];
				m_inputStrides[i] = m_inputStrides[i + 1] * input_dims[i + 1];
				m_inputStrides[i + 1] *= op.strides()[i + 1];
			}
			m_inputStrides[0] *= op.strides()[0];
		}
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dimensions; }

	EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType /*data*/)
	{
		m_impl.evalSubExprsIfNeeded(NULL);
		return true;
	}
	EIGEN_STRONG_INLINE void cleanup() { m_impl.cleanup(); }

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
	{
		return m_impl.coeff(srcCoeff(index));
	}

	template<int LoadMode>
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const
	{
		EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
		eigen_assert(index + PacketSize - 1 < dimensions().TotalSize());

		Index inputIndices[] = { 0, 0 };
		Index indices[] = { index, index + PacketSize - 1 };
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			EIGEN_UNROLL_LOOP
			for (int i = NumDims - 1; i > 0; --i) {
				const Index idx0 = indices[0] / m_outputStrides[i];
				const Index idx1 = indices[1] / m_outputStrides[i];
				inputIndices[0] += idx0 * m_inputStrides[i];
				inputIndices[1] += idx1 * m_inputStrides[i];
				indices[0] -= idx0 * m_outputStrides[i];
				indices[1] -= idx1 * m_outputStrides[i];
			}
			inputIndices[0] += indices[0] * m_inputStrides[0];
			inputIndices[1] += indices[1] * m_inputStrides[0];
		} else { // RowMajor
			EIGEN_UNROLL_LOOP
			for (int i = 0; i < NumDims - 1; ++i) {
				const Index idx0 = indices[0] / m_outputStrides[i];
				const Index idx1 = indices[1] / m_outputStrides[i];
				inputIndices[0] += idx0 * m_inputStrides[i];
				inputIndices[1] += idx1 * m_inputStrides[i];
				indices[0] -= idx0 * m_outputStrides[i];
				indices[1] -= idx1 * m_outputStrides[i];
			}
			inputIndices[0] += indices[0] * m_inputStrides[NumDims - 1];
			inputIndices[1] += indices[1] * m_inputStrides[NumDims - 1];
		}
		if (inputIndices[1] - inputIndices[0] == PacketSize - 1) {
			PacketReturnType rslt = m_impl.template packet<Unaligned>(inputIndices[0]);
			return rslt;
		} else {
			EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize];
			values[0] = m_impl.coeff(inputIndices[0]);
			values[PacketSize - 1] = m_impl.coeff(inputIndices[1]);
			EIGEN_UNROLL_LOOP
			for (int i = 1; i < PacketSize - 1; ++i) {
				values[i] = coeff(index + i);
			}
			PacketReturnType rslt = internal::pload<PacketReturnType>(values);
			return rslt;
		}
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const
	{
		double compute_cost = (NumDims - 1) * (TensorOpCost::AddCost<Index>() + TensorOpCost::MulCost<Index>() +
											   TensorOpCost::DivCost<Index>()) +
							  TensorOpCost::MulCost<Index>();
		if (vectorized) {
			compute_cost *= 2; // packet() computes two indices
		}
		const int innerDim = (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? 0 : (NumDims - 1);
		return m_impl.costPerCoeff(vectorized && m_inputStrides[innerDim] == 1) +
			   // Computation is not vectorized per se, but it is done once per packet.
			   TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
	}

	EIGEN_DEVICE_FUNC typename Storage::Type data() const { return NULL; }

#ifdef EIGEN_USE_SYCL
	// binding placeholder accessors to a command group handler for SYCL
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler& cgh) const { m_impl.bind(cgh); }
#endif
  protected:
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const
	{
		Index inputIndex = 0;
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			EIGEN_UNROLL_LOOP
			for (int i = NumDims - 1; i > 0; --i) {
				const Index idx = index / m_outputStrides[i];
				inputIndex += idx * m_inputStrides[i];
				index -= idx * m_outputStrides[i];
			}
			inputIndex += index * m_inputStrides[0];
		} else { // RowMajor
			EIGEN_UNROLL_LOOP
			for (int i = 0; i < NumDims - 1; ++i) {
				const Index idx = index / m_outputStrides[i];
				inputIndex += idx * m_inputStrides[i];
				index -= idx * m_outputStrides[i];
			}
			inputIndex += index * m_inputStrides[NumDims - 1];
		}
		return inputIndex;
	}

	Dimensions m_dimensions;
	array<Index, NumDims> m_outputStrides;
	array<Index, NumDims> m_inputStrides;
	TensorEvaluator<ArgType, Device> m_impl;
};

// Eval as lvalue
template<typename Strides, typename ArgType, typename Device>
struct TensorEvaluator<TensorStridingOp<Strides, ArgType>, Device>
	: public TensorEvaluator<const TensorStridingOp<Strides, ArgType>, Device>
{
	typedef TensorStridingOp<Strides, ArgType> XprType;
	typedef TensorEvaluator<const XprType, Device> Base;
	//  typedef typename XprType::Index Index;
	static const int NumDims = internal::array_size<typename TensorEvaluator<ArgType, Device>::Dimensions>::value;
	//  typedef DSizes<Index, NumDims> Dimensions;

	enum
	{
		IsAligned = /*TensorEvaluator<ArgType, Device>::IsAligned*/ false,
		PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
		PreferBlockAccess = false,
		Layout = TensorEvaluator<ArgType, Device>::Layout,
		CoordAccess = false, // to be implemented
		RawAccess = false
	};

	EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
		: Base(op, device)
	{
	}

	typedef typename XprType::Index Index;
	typedef typename XprType::Scalar Scalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
	static const int PacketSize = PacketType<CoeffReturnType, Device>::size;

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index)
	{
		return this->m_impl.coeffRef(this->srcCoeff(index));
	}

	template<int StoreMode>
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void writePacket(Index index, const PacketReturnType& x)
	{
		EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
		eigen_assert(index + PacketSize - 1 < this->dimensions().TotalSize());

		Index inputIndices[] = { 0, 0 };
		Index indices[] = { index, index + PacketSize - 1 };
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			EIGEN_UNROLL_LOOP
			for (int i = NumDims - 1; i > 0; --i) {
				const Index idx0 = indices[0] / this->m_outputStrides[i];
				const Index idx1 = indices[1] / this->m_outputStrides[i];
				inputIndices[0] += idx0 * this->m_inputStrides[i];
				inputIndices[1] += idx1 * this->m_inputStrides[i];
				indices[0] -= idx0 * this->m_outputStrides[i];
				indices[1] -= idx1 * this->m_outputStrides[i];
			}
			inputIndices[0] += indices[0] * this->m_inputStrides[0];
			inputIndices[1] += indices[1] * this->m_inputStrides[0];
		} else { // RowMajor
			EIGEN_UNROLL_LOOP
			for (int i = 0; i < NumDims - 1; ++i) {
				const Index idx0 = indices[0] / this->m_outputStrides[i];
				const Index idx1 = indices[1] / this->m_outputStrides[i];
				inputIndices[0] += idx0 * this->m_inputStrides[i];
				inputIndices[1] += idx1 * this->m_inputStrides[i];
				indices[0] -= idx0 * this->m_outputStrides[i];
				indices[1] -= idx1 * this->m_outputStrides[i];
			}
			inputIndices[0] += indices[0] * this->m_inputStrides[NumDims - 1];
			inputIndices[1] += indices[1] * this->m_inputStrides[NumDims - 1];
		}
		if (inputIndices[1] - inputIndices[0] == PacketSize - 1) {
			this->m_impl.template writePacket<Unaligned>(inputIndices[0], x);
		} else {
			EIGEN_ALIGN_MAX Scalar values[PacketSize];
			internal::pstore<Scalar, PacketReturnType>(values, x);
			this->m_impl.coeffRef(inputIndices[0]) = values[0];
			this->m_impl.coeffRef(inputIndices[1]) = values[PacketSize - 1];
			EIGEN_UNROLL_LOOP
			for (int i = 1; i < PacketSize - 1; ++i) {
				this->coeffRef(index + i) = values[i];
			}
		}
	}
};

} // end namespace Eigen

#endif // EIGEN_CXX11_TENSOR_TENSOR_STRIDING_H
